#include <errno.h>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <sys/ioctl.h>
#include "cutil_math.h"
#include "objects.h"
#include "ppm.h"
#include "scene.h"
#include "trace.h"

void cudaRaytrace(scene *c)
{
	
	/* load persistent data into constant memory */
	cudaMemcpyToSymbol(sfs, (void *) c->sf, sizeof(surface) * c->n_sfs, 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(out, (void *) &c->c.out, sizeof(size_t), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(ms, (void *) &c->m, sizeof(size_t), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(ts, (void *) &c->t, sizeof(size_t), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(ss, (void *) &c->s, sizeof(size_t), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(ps, (void *) &c->p, sizeof(size_t), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(ls, (void *) &c->l, sizeof(size_t), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(refract_idx, (void *) &c->refract_idx, sizeof(float), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(px_size, (void *) &c->c.px_size, sizeof(float), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(cam_mat, (void *) &c->c.cam_mat, 4 * sizeof(float4), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(max_bounces, (void *) &c->c.max_bounces, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(n_sfs, (void *) &c->n_sfs, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(n_ms, (void *) &c->n_ms, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(n_ts, (void *) &c->n_ts, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(n_ss, (void *) &c->n_ss, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(n_ps, (void *) &c->n_ps, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(n_ls, (void *) &c->n_ls, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(aa2, (void *) &c->c.aa2, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(aa, (void *) &c->c.aa, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(w, (void *) &c->c.x_res, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(h, (void *) &c->c.y_res, sizeof(int), 0, cudaMemcpyHostToDevice);
	
	dim3 block(BLOCK_W, BLOCK_H, 1);
	dim3 grid(c->grid_w, c->grid_h, 1);
	float fraction = 0;
	int total_pixels = c->c.x_res * c->c.y_res;
	size_t shared_mem = c->n_ms * sizeof(float);
	for(int i = 0; i < total_pixels; i += c->grid_w * c->grid_h)
	{
		// print status message an bar
		printf("\r|");
		for(int j = 0, l = ceil(fraction * (TERM_WIDTH - 3)); j < l; j++)
			putchar('-');
		putchar('>');
		for(int j = 0, l = floor((TERM_WIDTH - 3) - fraction * (TERM_WIDTH - 3)); j < l; j++)
			putchar(' ');
		printf("|\n");
		printf("%5.2f%% complete", 100 * fraction);
		fflush(stdout);
		
		
		doCudaRaytrace<<<grid, block, shared_mem>>>(i);
		fraction =  (float) i / total_pixels;
		
		// wait for block to finish
		cudaThreadSynchronize();
		
		// move cursor up one line
		printf("%c[1A", ASCII_ESC);
	}
	printf("\n                           \r");
	printf("100%% complete\n");
}

__global__
void doCudaRaytrace(int c_start)
{
	int thread_id = threadIdx.x + threadIdx.y * blockDim.x;
	if(thread_id == 0)
	{
		// x and y are the pixel coordinates
		x = (c_start + blockIdx.x + blockIdx.y * gridDim.x) % w;
		y = (c_start + blockIdx.x + blockIdx.y * gridDim.x) / w;
		
		// pixel number, from 0 to w * h - 1
		idx = x + y * w;
		
		// initialize output color to black
		px_color = make_float3(0);
		
		// only calculate the camera's starting position once
		camera_start = make_float3(cam_mat[0].w, cam_mat[1].w, cam_mat[2].w);
	}
	__syncthreads();
	
	// make sure the pixel actually exists
	if(idx >= w * h )
		return;
	
	// get the refraction of the starting point and store it in start_refract
	refract_start = getRefract(thread_id, camera_start);
	
	// loop across antialiasing iterations
	for(int i = 0; i < aa2; i++)
	{
		if(thread_id == 0)
		{
			// perct of light, from semi-reflective and semi-transparent surfaces
			reflectivity = 1.0;
			
			// refractive index of the current substance, assume it's air
			cur_refract = refract_start;
			
			// number of bounces
			bounce = max_bounces;
			
			// the ray direction slightly changes each time for antialiasing
			float4 tmp = make_float4(px_size *
			                        (x - w/2 +          (float) (i % aa) / aa),
			                        px_size *
			                        (h - y - 1 - h/2 +  (float) (i / aa) / aa),
			                        -1.0, 0.0);
			
			// set ray starting point and direction
			r.point = camera_start;
			r.direction.x = dot(cam_mat[0], tmp);
			r.direction.y = dot(cam_mat[1], tmp);
			r.direction.z = dot(cam_mat[2], tmp);
			r.direction = normalize(r.direction);
		}
		__syncthreads();
		
		// do the raytracing
		cudaTrace(thread_id);
	}
	
	// convert to rgb and write the pixel to global memory
	if(thread_id == 0)
	{
		out[idx].r = (unsigned char) min(255, 255 * px_color.x / aa2);
		out[idx].g = (unsigned char) min(255, 255 * px_color.y / aa2);
		out[idx].b = (unsigned char) min(255, 255 * px_color.z / aa2);
	}
}

__device__
void cudaTrace(int thread_id)
{
	do
	{
		// find the nearest intersection, and store it in shared memory
		getNearest(thread_id);
		if(thread_id == 0)
		{
			// change to our lights new position
			r.point += r.direction * distance;
			
			// type will be set to the number of triangles in the mesh
			if(type >= M_TRIANGLE)
			{
				// get the mesh triangle
				tmp_mt = getMTriangle((float *) object, type);
				
				// use phong interpolation
				float3 m_intersect = r.point - tmp_mt.point;
				float  m_x = dot(m_intersect, cross(tmp_mt.normal, tmp_mt.v[1])) /
				             dot(tmp_mt.v[0], cross(tmp_mt.normal, tmp_mt.v[1]));
				float m_y = dot(m_intersect, cross(tmp_mt.normal, tmp_mt.v[0])) /
				            dot(tmp_mt.v[1], cross(tmp_mt.normal, tmp_mt.v[0]));
				
				normal = normalize(tmp_mt.normals[0] * (1 - m_x - m_y) +
				                   tmp_mt.normals[1] * m_x +
				                   tmp_mt.normals[2] * m_y);
				
				// store the surface id
				sf_id = tmp_mt.sf_id;
				miss = 0;
			}
			else if(type == TRIANGLE)
			{
				tmp_t = getTriangle((float *) object, n_ts);
				normal = tmp_t.normal;
				sf_id = tmp_t.sf_id;
				miss = 0;
			}
			else if(type == SPHERE)
			{
				tmp_s = getSphere((float *) object, n_ss);
				normal = normalize(r.point - tmp_s.center);
				sf_id = tmp_s.sf_id;
				miss = 0;
			}
			else if(type == PLANE)
			{
				tmp_p = getPlane((float *) object, n_ps);
				normal = tmp_p.normal;
				sf_id = tmp_p.sf_id;
				miss = 0;
			}
			else
			{
				miss = 1;
				goto CHECK_BREAK; // horride, i know
			}
			
			// make sure the normal is pointing the right way
			if(dot(r.direction, normal) < 0)
				normal *= -1;
			
			// set the direction the light came from
			view_dir = -r.direction;
			
			// set the to-light vector's point
			to_l.point = r.point + EPSILON * view_dir;
			
			// if the object is opaque, reflect
			if(sfs[sf_id].opacity >= 1.0)
			{
				r.direction = reflect(r.direction, normal);
				
				// move our point off the surface to prevent hitting ourselves
				r.point += EPSILON * r.direction;
			}
			// refract
			else
			{
				float theta  = acos(dot(r.direction, normal));
				r.direction = (cur_refract / sfs[sf_id].refract) * r.direction +
				              normal * ((cur_refract / sfs[sf_id].refract) * cos(theta) -
				                        sqrt(1 - (sin(theta) * sin(theta))));
				r.direction = normalize(r.direction);
				r.point += EPSILON * r.direction;
			}
			
			
			// tmp color is this bounces color, initalize it
			tmp_color = make_float3(0.0);
			
			// set light counter, shared global for iterating across lights
			lt_counter = n_ls;
		}
		
		// we missed all objects, break out of the loop
		CHECK_BREAK:
		__syncthreads();
		if(miss)
			break;
		
		// get the new refraction
		cur_refract = getRefract(thread_id, r.point);
		__syncthreads();
		
		// iterate accros the lights
		while(lt_counter)
		{
			if(thread_id == 0)
			{
				lt_counter--;
				tmp_color += sfs[sf_id].ambient;
				
				// the the to-light vector's direction
				to_l.direction = normalize(ls[lt_counter].point - to_l.point);
				
				// a shared condition variable used in shadelight, set it now
				cv = 1.0;
				
				// shared variables used in shadelight, set them now
				shared_color_mult = make_float3(1.0, 1.0, 1.0);
				shared_transparency = 1.0;
			}
			__syncthreads();
			shadeLight(thread_id);
		}
		
		// factor in reflectivity/opacity and update overall reflectivity
		if(thread_id == 0)
		{
			if(sfs[sf_id].opacity < 1.0)
			{
				tmp_color *= sfs[sf_id].opacity;
				px_color += (tmp_color * reflectivity);
				reflectivity *= (1 - sfs[sf_id].opacity);
			}
			else
			{
				tmp_color *= (1 - sfs[sf_id].reflect);
				px_color += (tmp_color * reflectivity);
				reflectivity *= sfs[sf_id].reflect;
			}
			bounce--;
		}
		__syncthreads();
	} while(bounce > -1);
}

__device__
void getNearest(int thread_id)
{
	if(thread_id == 0)
	{
		distance = MAX_DIST;
		type = NONE;
	}
	
	// test all the free triangles
	if(n_ts > 0)
	{
		testTriangles(thread_id, ts, n_ts);
		if(thread_id == 0 && results_buffer[0] < distance)
		{
			distance = results_buffer[0];
			object = addrs_buffer[0];
			type = TRIANGLE;
			__threadfence_block();
		}
	}
	
	// test all the spheres
	if(n_ss > 0)
	{
		testSpheres(thread_id, ss, n_ss);
		if(thread_id == 0 && results_buffer[0] < distance)
		{
			distance = results_buffer[0];
			object = addrs_buffer[0];
			type = SPHERE;
			__threadfence_block();
		}
	}
	
	// test all the planes
	if(n_ps > 0)
	{
		testPlanes(thread_id, ps, n_ps);
		if(thread_id == 0 && results_buffer[0] < distance)
		{
			distance = results_buffer[0];
			object = addrs_buffer[0];
			type = PLANE;
			__threadfence_block();
		}
	}
	
	// test all the meshes
	if(n_ms > 0)
		testMeshes(thread_id, ms, n_ms);
}

__device__
void testTriangles(int thread_id, triangle *tris, int num_tris)
{
	float result = MAX_DIST;
	float tmp;
	void *p;
	
	// iterate and store the closest triangle's distance and address
	for(int i = thread_id; i < num_tris; i += NUM_THREADS)
	{
		tmp = triangleTest(getTriangle(((float *) tris) + i,  num_tris), r);
		if(tmp < result)
		{
			result = tmp;
			p = (void *) (((float*) tris) + i);
		}
	}
	
	// write results to the buffers and filter them
	results_buffer[thread_id] = result;
	addrs_buffer[thread_id] = p;
	__syncthreads();
	filterBuffers(thread_id);
}

__device__
void testSpheres(int thread_id, sphere *sphs, int num_sphs)
{
	float result = MAX_DIST;
	float tmp;
	void *p;
	
	// iterate and store the closest sphere's distance and address
	for(int i = thread_id; i < num_sphs; i = i + NUM_THREADS)
	{
		tmp = sphereTest(getSphere(((float *) sphs) + i, num_sphs), r);
		if(tmp < result)
		{
			result = tmp;
			p = (void *) (((float*) sphs) + i);
		}
	}
	
	// write results to the buffers and filter them
	results_buffer[thread_id] = result;
	addrs_buffer[thread_id] = p;
	__syncthreads();
	filterBuffers(thread_id);
}

__device__
void testPlanes(int thread_id, plane *plns, int num_plns)
{
	float result = MAX_DIST;
	float tmp;
	void *p;
	
	// iterate and store the closest plane's distance and address
	for(int i = thread_id; i < num_plns; i = i + NUM_THREADS)
	{
		tmp = planeTest(getPlane(((float *) plns) + i, num_plns), r);
		if(tmp < result)
		{
			result = tmp;
			p = (void *) (((float*) plns) + i);
		}
	}
	
	// write results to the buffers and filter them
	results_buffer[thread_id] = result;
	addrs_buffer[thread_id] = p;
	__syncthreads();
	filterBuffers(thread_id);
}

__device__
void testMeshes(int thread_id, mesh *mshs, int num_mshs)
{
	
	// store all the distances to the meshes' boxes in a shared buffer
	for(int i = thread_id; i < num_mshs; i = i + NUM_THREADS)
		mesh_buffer[i] = boxTest(getBox(((float *) mshs) + i, num_mshs), r);
	
	__syncthreads();
	
	// iterate across mesh buffer
	for(int i = 0; i < num_mshs; i++)
	{
		// only check meshes that are in front of the nearest thing
		if(mesh_buffer[i] < distance)
		{
			if(thread_id == 0)
				tmp_m = getLMesh(((float *) mshs) + i, num_mshs);
			__syncthreads();
			
			testTriangles(thread_id, (triangle *) tmp_m.t, tmp_m.n_ts);
			
			__syncthreads();
			
			// if we found a closer triangle, store it
			if(thread_id == 0 && results_buffer[0] < distance)
			{
				distance = results_buffer[0];
				object = addrs_buffer[0];
				
				// a hackish by free way of passing the number of triangles
				type = tmp_m.n_ts;
			}
		}
	}
	__syncthreads();
}

__device__
float m_triangleTest(m_triangle t, ray ry)
{
	if(!dot(ry.direction, t.normal))
		return MAX_DIST;
	
	// find plane
	float d = dot(t.normal, t.point);
	
	// solve for intersection point
	float distance = (dot(ry.point, t.normal) - d) /
	                 (-dot(ry.direction, t.normal));
	
	float3 intersect = ry.point +  ry.direction * distance - t.point;

	float x = dot(intersect, cross(t.normal, t.v[1])) /
	          dot(t.v[0],    cross(t.normal, t.v[1]));
	float y = dot(intersect, cross(t.normal, t.v[0])) /
	          dot(t.v[1],    cross(t.normal, t.v[0]));
	if(distance <= 0 || x <= 0 || y <= 0 || x + y > 1)
		return MAX_DIST;
	
	return distance;
}

__device__
float triangleTest(triangle t, ray ry)
{
	if(!dot(ry.direction, t.normal))
		return MAX_DIST;
	
	// find plane
	float d = dot(t.normal, t.point);
	
	// solve for intersection point
	float distance = (dot(ry.point, t.normal) - d) /
	                 (-dot(ry.direction, t.normal));
	
	float3 intersect = ry.point +  ry.direction * distance - t.point;

	float x = dot(intersect, cross(t.normal, t.v[1])) /
	          dot(t.v[0],    cross(t.normal, t.v[1]));
	float y = dot(intersect, cross(t.normal, t.v[0])) /
	          dot(t.v[1],    cross(t.normal, t.v[0]));
	if(distance <= 0 || x <= 0 || y <= 0 || x + y > 1)
		return MAX_DIST;
	
	return distance;
}


__device__
float planeTest(plane p, ray ry)
{
	if(!dot(ry.direction, p.normal))
		return MAX_DIST;
	
	float d = dot(p.normal, p.point);
	
	float t = (dot(ry.point, p.normal) - d) /
	          -dot(ry.direction, p.normal);
	if(t > 0)
		return t;
	
	return MAX_DIST;
}

__device__
float sphereTest(sphere s, ray ry)
{
	float a = dot(ry.direction, ry.direction);
	float b = 2 * dot(ry.direction, ry.point - s.center);
	float c = dot(ry.point - s.center, ry.point - s.center) -
	          s.radius * s.radius;
	float discriminant = b * b - 4 * a * c;
	if(discriminant < 0)
		return MAX_DIST;
	
	float answer = (-b - sqrt(discriminant)) / (2 * a);
	
	if(answer <= 0)
		return MAX_DIST;
	
	return answer;
}


/*
 * Check whether we hit a mesh's box. It's just a super simple way to
 * avoid some unecessary work. We check if the ray hits the mesh's box
 * before we test its triangles.
 */
__device__
float boxTest(box b, ray ry)
{
	// if the ray starts inside the box, we always have to check its mesh
	if(ry.point.x > b.mins.x && ry.point.x < b.maxs.x &&
	   ry.point.y > b.mins.y && ry.point.y < b.maxs.y &&
	   ry.point.z > b.mins.z && ry.point.z < b.maxs.z)
		return 0;
	
	// nvcc puts arrays in local memory, so no array
	float intersect1;
	float intersect2;
	float tmp;
	
	// check low x face
	tmp = b.mins.x - ry.point.x;
	if(tmp > 0 && ry.direction.x > 0)
	{
		tmp /= r.direction.x;
		intersect1 = tmp * ry.direction.y + ry.point.y;
		intersect2 = tmp * ry.direction.z + ry.point.z;
		if(intersect1 >= b.mins.y && intersect1 <= b.maxs.y &&
		   intersect2 >= b.mins.z && intersect2 <= b.maxs.z);
			return tmp;
	}
	
	// high x face
	tmp = ry.point.x - b.maxs.x;
	if(tmp > 0 && ry.direction.x < 0)
	{
		tmp /= r.direction.x;
		intersect1 = tmp * ry.direction.y + ry.point.y;
		intersect2 = tmp * ry.direction.z + ry.point.z;
		if(intersect1 >= b.mins.y && intersect1 <= b.maxs.y &&
		   intersect2 >= b.mins.z && intersect2 <= b.maxs.z);
			return tmp;
	}
	
	// low y face
	tmp = b.mins.y - ry.point.y;
	if(tmp > 0 && ry.direction.y > 0)
	{
		tmp /= r.direction.y;
		intersect1 = tmp * ry.direction.x + ry.point.x;
		intersect2 = tmp * ry.direction.z + ry.point.z;
		if(intersect1 >= b.mins.x && intersect1 <= b.maxs.x &&
		   intersect2 >= b.mins.z && intersect2 <= b.maxs.z);
			return tmp;
	}
	
	// high y face
	tmp = ry.point.y - b.maxs.y;
	if(tmp > 0 && ry.direction.y < 0)
	{
		tmp /= r.direction.y;
		intersect1 = tmp * ry.direction.x + ry.point.x;
		intersect2 = tmp * ry.direction.z + ry.point.z;
		if(intersect1 >= b.mins.x && intersect1 <= b.maxs.x &&
		   intersect2 >= b.mins.z && intersect2 <= b.maxs.z);
			return tmp;
	}
	
	// low z face
	tmp = b.mins.z - ry.point.z;
	if(tmp > 0 && ry.direction.z > 0)
	{
		tmp /= r.direction.z;
		intersect1 = tmp * ry.direction.x + ry.point.x;
		intersect2 = tmp * ry.direction.y + ry.point.y;
		if(intersect1 >= b.mins.x && intersect1 <= b.maxs.x &&
		   intersect2 >= b.mins.y && intersect2 <= b.maxs.y);
			return tmp;
	}
	
	// high z face
	tmp = ry.point.z - b.maxs.z;
	if(tmp > 0 && ry.direction.z < 0)
	{
		tmp /= r.direction.z;
		intersect1 = tmp * ry.direction.x + ry.point.x;
		intersect2 = tmp * ry.direction.y + ry.point.y;
		if(intersect1 >= b.mins.x && intersect1 <= b.maxs.x &&
		   intersect2 >= b.mins.y && intersect2 <= b.maxs.y);
			return tmp;
	}
	
	// we didn't hit any faces
	return MAX_DIST;
}


/*
 * Filter through the results buffer to find the nearest object. Copy the
 * object pointers too.
 *
 * One warp goes across the buffers and finds the nearest for each thread,
 * and the thread 0 finds the best of those.
 */

__device__
void filterBuffers(int thread_id)
{
	float result;
	void *p;
	if(thread_id < warpSize)
	{
		result = MAX_DIST;
		for(int i = thread_id; i < BLOCK_W * BLOCK_H; i += warpSize)
		{
			if(results_buffer[i] < result)
			{
				result = results_buffer[i];
				p      = addrs_buffer[i];
			}
		}
		results_buffer[thread_id] = result;
		addrs_buffer[thread_id]   = p;
		__threadfence_block();
	}
	
	__syncthreads();
	if(thread_id == 0)
	{
		result = MAX_DIST;
		for(int i= 0; i < warpSize; i++)
			if(results_buffer[i] < result)
			{
				result = results_buffer[i];
				p      = addrs_buffer[i];
			}
		results_buffer[0] = result;
		addrs_buffer[0]   = p;
	}
	__syncthreads();
}


/*
 * Get color contribution from the current light. If the object is in
 * shadow, don't contribute anything, and tell all the threads to stop
 * checking objects. If it is behind something that is semi-transparent,
 * dim as necessary, and use its diffuse color to shade the light that
 * passes through.
 */
__device__
void shadeLight(int thread_id)
{
	// start out the color multiplier as white
	float3 color_mult = make_float3(1.0, 1.0, 1.0);
	float transparency = 1.0;
	// check all the triangles
	for(int i = thread_id; i < n_ts && cv; i += NUM_THREADS)
	{
		triangle t = getTriangle(((float *) ts) + i, n_ts);
		if(triangleTest(t, to_l) < MAX_DIST)
		{
			// chage the color multiplier based on the opacity and color
			if(sfs[t.sf_id].opacity < 1.0)
			{
				transparency *= 1 - sfs[t.sf_id].opacity;
				color_mult.x *= 1 - sfs[t.sf_id].opacity * (1 - sfs[t.sf_id].diffuse.x);
				color_mult.y *= 1 - sfs[t.sf_id].opacity * (1 - sfs[t.sf_id].diffuse.y);
				color_mult.z *= 1 - sfs[t.sf_id].opacity * (1 - sfs[t.sf_id].diffuse.z);
			}
			
			// if the object is opaque, stop everything
			else
			{
				cv = 0;
				__threadfence_block();
			}
		}
	}
	__syncthreads();
	if(!cv)
		return;
	
	// check all the spheres
	for(int i = thread_id; i < n_ss && cv; i += NUM_THREADS)
	{
		sphere s = getSphere(((float *) ss) + i, n_ss);
		if(sphereTest(s, to_l) < MAX_DIST)
		{
			// chage the color multiplier based on the opacity and color
			if(sfs[s.sf_id].opacity < 1.0)
			{
				transparency *= 1 - sfs[s.sf_id].opacity;
				color_mult.x *= 1 - sfs[s.sf_id].opacity * (1 - sfs[s.sf_id].diffuse.x);
				color_mult.y *= 1 - sfs[s.sf_id].opacity * (1 - sfs[s.sf_id].diffuse.y);
				color_mult.z *= 1 - sfs[s.sf_id].opacity * (1 - sfs[s.sf_id].diffuse.z);
			}
			
			// if the object is opaque, stop everything
			else
			{
				cv = 0;
				__threadfence_block();
			}
		}
	}
	__syncthreads();
	if(!cv)
		return;
	
	// check all the planes
	for(int i = thread_id; i < n_ps && cv; i += NUM_THREADS)
	{
		plane p = getPlane(((float *) ps) + i, n_ps);
		if(planeTest(p, to_l) < MAX_DIST)
		{
			// chage the color multiplier based on the opacity and color
			if(sfs[p.sf_id].opacity < 1.0)
			{
				transparency *= 1 - sfs[p.sf_id].opacity;
				color_mult.x *= 1 - sfs[p.sf_id].opacity * (1 - sfs[p.sf_id].diffuse.x);
				color_mult.y *= 1 - sfs[p.sf_id].opacity * (1 - sfs[p.sf_id].diffuse.y);
				color_mult.z *= 1 - sfs[p.sf_id].opacity * (1 - sfs[p.sf_id].diffuse.z);
			}
			
			// if the object is opaque, stop everything
			else
			{
				cv = 0;
				__threadfence_block();
			}
		}
	}
	__syncthreads();
	if(!cv)
		return;
	
	// check all the meshes
	for(int i = 0; i < n_ms && cv; i++)
	{
		// first check the mesh's box
		if(thread_id == 0)
		{
			if(boxTest(getBox(((float *) ms) + i, n_ms), to_l) < MAX_DIST)
				tmp_m = getLMesh(((float *) ms) + i, n_ms);
			
			// if we don't hit it, use the number of triangles as a signal to stop
			else
				tmp_m.n_ts = 0;
			
			__threadfence_block();
		}
		__syncthreads();
		if(tmp_m.n_ts == 0)
			continue;
		
		for(int j = thread_id; j < tmp_m.n_ts && cv; j += NUM_THREADS)
		{
			m_triangle t = getMTriangle(((float *) tmp_m.t) + j, tmp_m.n_ts);
			if(m_triangleTest(t, to_l) < MAX_DIST)
			{
				// chage the color multiplier based on the opacity and color
				if(sfs[t.sf_id].opacity < 1.0)
				{
					transparency *= 1 - sfs[t.sf_id].opacity;
					color_mult.x *= 1 - sfs[t.sf_id].opacity * (1 - sfs[t.sf_id].diffuse.x);
					color_mult.y *= 1 - sfs[t.sf_id].opacity * (1 - sfs[t.sf_id].diffuse.y);
					color_mult.z *= 1 - sfs[t.sf_id].opacity * (1 - sfs[t.sf_id].diffuse.z);
				}
				
				// if the object is opaque, stop everything
				else
				{
					cv = 0;
					__threadfence_block();
				}
			}
		}
	}
	
	__syncthreads();
	if(!cv)
		return;
	
	// combine all the color multipliers
	for(int i = 0; i < NUM_THREADS; i++)
		if(i == thread_id)
		{
			shared_color_mult.x *= color_mult.x;
			shared_color_mult.y *= color_mult.y;
			shared_color_mult.z *= color_mult.z;
			shared_transparency *= transparency;
			__threadfence_block();
		}
	
	if(thread_id == 0)
	{
		shared_color_mult *= shared_transparency;
		// bounce direction, used for calcuating specular lighting
		float3 b_dir = -1 * reflect(to_l.direction, normal);
		
		// we need abs for inverted mesh triangles
		float3 dif = sfs[sf_id].diffuse * ls[lt_counter].intensity *
		             max(0., abs(dot(to_l.direction, normal)));
		float3 spec = sfs[sf_id].specular * ls[lt_counter].intensity *
		              powf(max(0., dot(b_dir, view_dir)), sfs[sf_id].specular_exp);
		dif.x *= shared_color_mult.x;
		dif.y *= shared_color_mult.y;
		dif.z *= shared_color_mult.z;
		spec.x *= shared_color_mult.x;
		spec.y *= shared_color_mult.y;
		spec.z *= shared_color_mult.z;
		tmp_color += dif + spec;
		__threadfence_block();
	}
}

/*
 * Get the current refractive index. If were in a sphere, get its
 * index, otherwise use the default.
 */
__device__
float getRefract(int thread_id, float3 p)
{
	float result = MAX_DIST;
	float tmp;
	void *obj;
	sphere s;
	
	// check all the spheres to see if we're in one
	for(int i = thread_id; i < n_ss; i = i + NUM_THREADS)
	{
		s = getSphere(((float *) ss) + i, n_ss);
		tmp = length(p - s.center);
		
		// we only care about the sphere we iide of and closest to the edge of
		if(tmp < s.radius && tmp < result)
		{
			result = tmp;
			obj = ((void *) ((float *) ss + i));
		}
	}
	results_buffer[thread_id] = result;
	addrs_buffer[thread_id] = obj;
	__syncthreads();
	filterBuffers(thread_id);
	
	if(thread_id == 0)
	{
		// if we're not in a sphere, return the default refractive index
		if(results_buffer[0] == MAX_DIST)
			return refract_idx;
		
		// return the nearest spheres refractive index
		s = getSphere((float*) addrs_buffer[0], n_ss);
		return sfs[s.sf_id].refract;
	}
	
	// just so the compiler stops complaining
	return refract_idx;
}
